#include "kernel_operator.h"

using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
template<typename T>
class KernelAddcmul {
 public:
  __aicore__ inline KernelAddcmul() {}
  __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR value, GM_ADDR out,
                              uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t formerNum,
                              uint32_t formerLength, uint32_t formertileNum,
                              uint32_t formertileLength,
                              uint32_t formerlasttileLength, uint32_t tailNum,
                              uint32_t tailLength, uint32_t tailtileNum,
                              uint32_t tailtileLength,
                              uint32_t taillasttileLength, uint32_t tilingKey, uint32_t datatype) {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
    this->datatype = datatype;
    valueGm.SetGlobalBuffer((__gm__ T*)value , 1);
    this->value = valueGm.GetValue(0);

    if (tilingKey == 1) {
        this->blockLength = blockLength;
        this->tileNum =
            tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
        this->tileLength = tileLength / BUFFER_NUM;
        this->lasttileLength = lasttileLength;

        xGm.SetGlobalBuffer((__gm__ T*)x + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        yGm.SetGlobalBuffer((__gm__ T*)y + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        zGm.SetGlobalBuffer((__gm__ T*)z + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        // valueGm.SetGlobalBuffer((__gm__ T*)value + this->blockLength * GetBlockIdx(),
        //                     this->blockLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->blockLength * GetBlockIdx(),
            this->blockLength);
        }

    if (tilingKey == 2) {
      this->formerNum = formerNum;
      this->formerLength = formerLength;
      this->formertileNum = formertileNum;
      this->formertileLength = formertileLength;
      this->formerlasttileLength = formerlasttileLength;

      this->tailNum = tailNum;
      this->tailLength = tailLength;
      this->tailtileNum = tailtileNum;
      this->tailtileLength = tailtileLength;
      this->taillasttileLength = taillasttileLength;

      if (GetBlockIdx() < this->formerNum) {  //分到大块核的处理
        this->tileLength = this->formertileLength / BUFFER_NUM;
        this->lasttileLength = this->formerlasttileLength;
        this->tileNum = this->formertileNum * BUFFER_NUM;
        xGm.SetGlobalBuffer(
            (__gm__ T*)x + this->formerLength * GetBlockIdx(),
            this->formerLength);
        yGm.SetGlobalBuffer(
            (__gm__ T*)y + this->formerLength * GetBlockIdx(),
            this->formerLength);
        zGm.SetGlobalBuffer(
            (__gm__ T*)z + this->formerLength * GetBlockIdx(),
            this->formerLength);
        // valueGm.SetGlobalBuffer(
        //     (__gm__ T*)value + this->formerLength * GetBlockIdx(),
        //     this->formerLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->formerLength * GetBlockIdx(),
            this->formerLength);
      } else {  //分到小块核的处理，需要处理的数据量比大核少alignNum个
        this->tileLength = this->tailtileLength / BUFFER_NUM;
        this->lasttileLength = this->taillasttileLength;
        this->tileNum = this->tailtileNum * BUFFER_NUM;
        xGm.SetGlobalBuffer(
            (__gm__ T*)x + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        yGm.SetGlobalBuffer(
            (__gm__ T*)y + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        zGm.SetGlobalBuffer(
            (__gm__ T*)z + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        // valueGm.SetGlobalBuffer(
        //     (__gm__ T*)value + this->formerLength * this->formerNum +
        //         this->tailLength * (GetBlockIdx() - this->formerNum),
        //     this->tailLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
      }
    }
    
    pipe.InitBuffer(inQueueIN, BUFFER_NUM, this->tileLength * 4 * sizeof(T));
    pipe.InitBuffer(outQueueOUT, BUFFER_NUM, this->tileLength * sizeof(T));
  }
  __aicore__ inline void Process() {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++) {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }
  }

 private:
  __aicore__ inline void CopyIn(int32_t progress) {
    LocalTensor<T> inLocal = inQueueIN.AllocTensor<T>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
            //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
            DataCopy(inLocal[0], xGm[0], this->tileLength);
            DataCopy(inLocal[this->tileLength], yGm[0], this->tileLength);
            DataCopy(inLocal[2 * (this->tileLength)], zGm[0], this->tileLength);
            DataCopy(inLocal[3 * (this->tileLength)], zGm[0], this->tileLength);
        } else {
            //将最后一个分块的起始地址向前移动tileLength-lasttileLength
            DataCopy(
                inLocal[0],
                xGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);
            DataCopy(
                inLocal[this->tileLength],
                yGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);
            DataCopy(
                inLocal[2 * (this->tileLength)],
                zGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);
            // DataCopy(
            //     inLocal[3 * (this->tileLength)],
            //     valueGm[(progress - 1) * this->tileLength + this->lasttileLength],
            //     this->tileLength);
        }
      } else {
            DataCopy(inLocal[0], xGm[progress * this->tileLength],
                    this->tileLength);
            DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
                    this->tileLength);
            DataCopy(inLocal[2 * (this->tileLength)],
                    zGm[progress * this->tileLength], this->tileLength);
            // DataCopy(inLocal[3 * (this->tileLength)],
            //         valueGm[progress * this->tileLength], this->tileLength);
      }
    }
    if (BUFFER_NUM == 2) {
      //开启double
      //buffer时，由于将输入数据分成了相等的2部分，分块大小为不开启double
      //buffer的一半， 所以需要对最后两个分块数据的起始地址做处理
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        //分块大小变为tileLength的一半
        //倒数第2个分块数据的起始地址向前移动（tileLength-lasttileLength)，最后一个分块的起始地址以此为基础进行移动
        DataCopy(
            inLocal[0],
            xGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        DataCopy(
            inLocal[this->tileLength],
            yGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        DataCopy(
            inLocal[2 * (this->tileLength)],
            zGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        // DataCopy(
        //     inLocal[3 * (this->tileLength)],
        //     valueGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
        //     (this->tileLength));

      }

      else {
        DataCopy(inLocal[0], xGm[progress * (this->tileLength)],
                 (this->tileLength));
        DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
                 this->tileLength);
        DataCopy(inLocal[2 * (this->tileLength)],
                 zGm[progress * this->tileLength], this->tileLength);
        // DataCopy(inLocal[3 * (this->tileLength)],
        //          valueGm[progress * this->tileLength], this->tileLength);
      }
    }

    inQueueIN.EnQue(inLocal);
  }
  __aicore__ inline void Compute(int32_t progress) {
    LocalTensor<T> inLocal = inQueueIN.DeQue<T>();
    LocalTensor<T> xLocal = inLocal;
    LocalTensor<T> yLocal = inLocal[this->tileLength];
    LocalTensor<T> zLocal = inLocal[2 * (this->tileLength)];

    LocalTensor<T> outLocal = outQueueOUT.AllocTensor<T>();
    
    if(this->datatype == 2){
      TBuf<TPosition::VECCALC> xBuf;
      TBuf<TPosition::VECCALC> yBuf;
      TBuf<TPosition::VECCALC> zBuf;
      TBuf<TPosition::VECCALC> outBuf;
      TBuf<TPosition::VECCALC> valueBuf;
      pipe.InitBuffer(xBuf, this->tileLength * sizeof(half));
      pipe.InitBuffer(yBuf, this->tileLength * sizeof(half));
      pipe.InitBuffer(zBuf, this->tileLength * sizeof(half));
      pipe.InitBuffer(outBuf, this->tileLength * sizeof(half));
      pipe.InitBuffer(valueBuf, this->tileLength * sizeof(half));
      
      LocalTensor<half> xTemp = xBuf.Get<half>();
      LocalTensor<half> yTemp = yBuf.Get<half>();
      LocalTensor<half> zTemp = zBuf.Get<half>();
      LocalTensor<half> outTemp = outBuf.Get<half>();
      LocalTensor<half> valueTemp = valueBuf.Get<half>();
      half temp = (half)this->value;

      Cast(xTemp, xLocal, RoundMode::CAST_NONE, this->tileLength);
      Cast(yTemp, yLocal, RoundMode::CAST_NONE, this->tileLength);
      Cast(zTemp, zLocal, RoundMode::CAST_NONE, this->tileLength);
      Cast(outTemp, outLocal, RoundMode::CAST_NONE, this->tileLength);

      Mul(outTemp, yTemp, zTemp, this->tileLength);
      Muls(outTemp, outTemp, temp, this->tileLength);
      Add(outTemp, xTemp, outTemp, this->tileLength);

      Cast(outLocal, outTemp, RoundMode::CAST_NONE, this->tileLength);
  } 
  else{
    Mul(outLocal, yLocal, zLocal, this->tileLength);
    Muls(outLocal, outLocal, this->value, this->tileLength);
    Add(outLocal, xLocal, outLocal, this->tileLength);
  }
    
    outQueueOUT.EnQue<T>(outLocal);

    inQueueIN.FreeTensor(inLocal);
  }
  __aicore__ inline void CopyOut(int32_t progress) {
    LocalTensor<T> outLocal = outQueueOUT.DeQue<T>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
          //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
          DataCopy(outGm[0], outLocal, this->tileLength);
        } else {
          //将最后一个分块的起始地址向前移动tileLength-lasttileLength
          DataCopy(
              outGm[(progress - 1) * this->tileLength + this->lasttileLength],
              outLocal, this->tileLength);
        }
      } else {
        DataCopy(outGm[progress * this->tileLength], outLocal,
                 this->tileLength);
      }
    }
    if (BUFFER_NUM == 2) {
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        DataCopy(
            outGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            outLocal, (this->tileLength));
      }

      else {
        DataCopy(outGm[progress * (this->tileLength)], outLocal,
                 (this->tileLength));
      }
    }

    outQueueOUT.FreeTensor(outLocal);
  }

 private:
  TPipe pipe;
  // TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueZ;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
  GlobalTensor<T> xGm;
  GlobalTensor<T> yGm;
  GlobalTensor<T> zGm;
  GlobalTensor<T> valueGm;
  GlobalTensor<T> outGm;
  T value;
  uint32_t datatype;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t formerNum;
  uint32_t formerLength;
  uint32_t formertileNum;
  uint32_t formertileLength;
  uint32_t formerlasttileLength;
  uint32_t tailNum;
  uint32_t tailLength;
  uint32_t tailtileNum;
  uint32_t tailtileLength;
  uint32_t taillasttileLength;
};

template<typename T>
class KernelAddcmul_for_half {
 public:
  __aicore__ inline KernelAddcmul_for_half() {}
  __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR value, GM_ADDR out,
                              uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t formerNum,
                              uint32_t formerLength, uint32_t formertileNum,
                              uint32_t formertileLength,
                              uint32_t formerlasttileLength, uint32_t tailNum,
                              uint32_t tailLength, uint32_t tailtileNum,
                              uint32_t tailtileLength,
                              uint32_t taillasttileLength, uint32_t tilingKey, uint32_t datatype) {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
    this->datatype = datatype;
    //valueGm.SetGlobalBuffer((__gm__ T*)value , 1);
    //this->value = valueGm.GetValue(0);

    if (tilingKey == 1) {
        this->blockLength = blockLength;
        this->tileNum =
            tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
        this->tileLength = tileLength / BUFFER_NUM;
        this->lasttileLength = lasttileLength;

        xGm.SetGlobalBuffer((__gm__ T*)x + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        yGm.SetGlobalBuffer((__gm__ T*)y + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        zGm.SetGlobalBuffer((__gm__ T*)z + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        valueGm.SetGlobalBuffer((__gm__ T*)value + this->blockLength * GetBlockIdx(),
                            this->blockLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->blockLength * GetBlockIdx(),
            this->blockLength);
        }

    if (tilingKey == 2) {
      this->formerNum = formerNum;
      this->formerLength = formerLength;
      this->formertileNum = formertileNum;
      this->formertileLength = formertileLength;
      this->formerlasttileLength = formerlasttileLength;

      this->tailNum = tailNum;
      this->tailLength = tailLength;
      this->tailtileNum = tailtileNum;
      this->tailtileLength = tailtileLength;
      this->taillasttileLength = taillasttileLength;

      if (GetBlockIdx() < this->formerNum) {  //分到大块核的处理
        this->tileLength = this->formertileLength / BUFFER_NUM;
        this->lasttileLength = this->formerlasttileLength;
        this->tileNum = this->formertileNum * BUFFER_NUM;
        xGm.SetGlobalBuffer(
            (__gm__ T*)x + this->formerLength * GetBlockIdx(),
            this->formerLength);
        yGm.SetGlobalBuffer(
            (__gm__ T*)y + this->formerLength * GetBlockIdx(),
            this->formerLength);
        zGm.SetGlobalBuffer(
            (__gm__ T*)z + this->formerLength * GetBlockIdx(),
            this->formerLength);
        valueGm.SetGlobalBuffer(
            (__gm__ T*)value + this->formerLength * GetBlockIdx(),
            this->formerLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->formerLength * GetBlockIdx(),
            this->formerLength);
      } else {  //分到小块核的处理，需要处理的数据量比大核少alignNum个
        this->tileLength = this->tailtileLength / BUFFER_NUM;
        this->lasttileLength = this->taillasttileLength;
        this->tileNum = this->tailtileNum * BUFFER_NUM;
        xGm.SetGlobalBuffer(
            (__gm__ T*)x + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        yGm.SetGlobalBuffer(
            (__gm__ T*)y + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        zGm.SetGlobalBuffer(
            (__gm__ T*)z + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        valueGm.SetGlobalBuffer(
            (__gm__ T*)value + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
        outGm.SetGlobalBuffer(
            (__gm__ T*)out + this->formerLength * this->formerNum +
                this->tailLength * (GetBlockIdx() - this->formerNum),
            this->tailLength);
      }
    }
    this->value = valueGm.GetValue(0);
    pipe.InitBuffer(inQueueIN, BUFFER_NUM, this->tileLength * 4 * sizeof(T));
    pipe.InitBuffer(outQueueOUT, BUFFER_NUM, this->tileLength * sizeof(T));
  }
  __aicore__ inline void Process() {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++) {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }
  }

 private:
  __aicore__ inline void CopyIn(int32_t progress) {
    LocalTensor<T> inLocal = inQueueIN.AllocTensor<T>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
            //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
            DataCopy(inLocal[0], xGm[0], this->tileLength);
            DataCopy(inLocal[this->tileLength], yGm[0], this->tileLength);
            DataCopy(inLocal[2 * (this->tileLength)], zGm[0], this->tileLength);
            DataCopy(inLocal[3 * (this->tileLength)], zGm[0], this->tileLength);
        } else {
            //将最后一个分块的起始地址向前移动tileLength-lasttileLength
            DataCopy(
                inLocal[0],
                xGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);
            DataCopy(
                inLocal[this->tileLength],
                yGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);
            DataCopy(
                inLocal[2 * (this->tileLength)],
                zGm[(progress - 1) * this->tileLength + this->lasttileLength],
                this->tileLength);

        }
      } else {
            DataCopy(inLocal[0], xGm[progress * this->tileLength],
                    this->tileLength);
            DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
                    this->tileLength);
            DataCopy(inLocal[2 * (this->tileLength)],
                    zGm[progress * this->tileLength], this->tileLength);
            // DataCopy(inLocal[3 * (this->tileLength)],
            //         valueGm[progress * this->tileLength], this->tileLength);
      }
    }
    if (BUFFER_NUM == 2) {
      //开启double
      //buffer时，由于将输入数据分成了相等的2部分，分块大小为不开启double
      //buffer的一半， 所以需要对最后两个分块数据的起始地址做处理
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        //分块大小变为tileLength的一半
        //倒数第2个分块数据的起始地址向前移动（tileLength-lasttileLength)，最后一个分块的起始地址以此为基础进行移动
        DataCopy(
            inLocal[0],
            xGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        DataCopy(
            inLocal[this->tileLength],
            yGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        DataCopy(
            inLocal[2 * (this->tileLength)],
            zGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));
        // DataCopy(
        //     inLocal[3 * (this->tileLength)],
        //     valueGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
        //     (this->tileLength));

      }

      else {
        DataCopy(inLocal[0], xGm[progress * (this->tileLength)],
                 (this->tileLength));
        DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
                 this->tileLength);
        DataCopy(inLocal[2 * (this->tileLength)],
                 zGm[progress * this->tileLength], this->tileLength);
        // DataCopy(inLocal[3 * (this->tileLength)],
        //          valueGm[progress * this->tileLength], this->tileLength);
      }
    }

    inQueueIN.EnQue(inLocal);
  }
  __aicore__ inline void Compute(int32_t progress) {
    LocalTensor<T> inLocal = inQueueIN.DeQue<T>();
    LocalTensor<T> xLocal = inLocal;
    LocalTensor<T> yLocal = inLocal[this->tileLength];
    LocalTensor<T> zLocal = inLocal[2 * (this->tileLength)];

    LocalTensor<T> outLocal = outQueueOUT.AllocTensor<T>();
  
    Mul(outLocal, yLocal, zLocal, this->tileLength);
    Muls(outLocal, outLocal, this->value, this->tileLength);
    Add(outLocal, xLocal, outLocal, this->tileLength);
  
    
    outQueueOUT.EnQue<T>(outLocal);

    inQueueIN.FreeTensor(inLocal);
  }
  __aicore__ inline void CopyOut(int32_t progress) {
    LocalTensor<T> outLocal = outQueueOUT.DeQue<T>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
          //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
          DataCopy(outGm[0], outLocal, this->tileLength);
        } else {
          //将最后一个分块的起始地址向前移动tileLength-lasttileLength
          DataCopy(
              outGm[(progress - 1) * this->tileLength + this->lasttileLength],
              outLocal, this->tileLength);
        }
      } else {
        DataCopy(outGm[progress * this->tileLength], outLocal,
                 this->tileLength);
      }
    }
    if (BUFFER_NUM == 2) {
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        DataCopy(
            outGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            outLocal, (this->tileLength));
      }

      else {
        DataCopy(outGm[progress * (this->tileLength)], outLocal,
                 (this->tileLength));
      }
    }

    outQueueOUT.FreeTensor(outLocal);
  }

 private:
  TPipe pipe;
  // TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueZ;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
  GlobalTensor<T> xGm;
  GlobalTensor<T> yGm;
  GlobalTensor<T> zGm;
  GlobalTensor<T> valueGm;
  GlobalTensor<T> outGm;
  T value;
  uint32_t datatype;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t formerNum;
  uint32_t formerLength;
  uint32_t formertileNum;
  uint32_t formertileLength;
  uint32_t formerlasttileLength;
  uint32_t tailNum;
  uint32_t tailLength;
  uint32_t tailtileNum;
  uint32_t tailtileLength;
  uint32_t taillasttileLength;
};


extern "C" __global__ __aicore__ void addcmul(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR value, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    if(tiling_data.datatype == 2){
      //int8
      KernelAddcmul<int8_t> op;

      uint32_t tilingKey = 1;
      if (TILING_KEY_IS(1)) {
          tilingKey = 1;
      } else if (TILING_KEY_IS(2)) {
          tilingKey = 2;
      } else {
          tilingKey = 1;
      }
      op.Init(x, y, z, value, out, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.formerNum,
            tiling_data.formerLength, tiling_data.formertileNum,
            tiling_data.formertileLength, tiling_data.formerlasttileLength,
            tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
            tiling_data.tailtileLength, tiling_data.taillasttileLength,
            tilingKey, tiling_data.datatype);
      op.Process();
    }else if(tiling_data.datatype == 0){
      //float
      KernelAddcmul<float> op;

      uint32_t tilingKey = 1;
      if (TILING_KEY_IS(1)) {
          tilingKey = 1;
      } else if (TILING_KEY_IS(2)) {
          tilingKey = 2;
      } else {
          tilingKey = 1;
      }
      op.Init(x, y, z, value, out, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.formerNum,
            tiling_data.formerLength, tiling_data.formertileNum,
            tiling_data.formertileLength, tiling_data.formerlasttileLength,
            tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
            tiling_data.tailtileLength, tiling_data.taillasttileLength,
            tilingKey, tiling_data.datatype);
      op.Process();
    }else if(tiling_data.datatype == 3){
      //int32
      KernelAddcmul<int32_t> op;

      uint32_t tilingKey = 1;
      if (TILING_KEY_IS(1)) {
          tilingKey = 1;
      } else if (TILING_KEY_IS(2)) {
          tilingKey = 2;
      } else {
          tilingKey = 1;
      }
      op.Init(x, y, z, value, out, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.formerNum,
            tiling_data.formerLength, tiling_data.formertileNum,
            tiling_data.formertileLength, tiling_data.formerlasttileLength,
            tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
            tiling_data.tailtileLength, tiling_data.taillasttileLength,
            tilingKey, tiling_data.datatype);
      op.Process();
    }
    else if(tiling_data.datatype == 1){
      //half
      KernelAddcmul_for_half<half> op;

      uint32_t tilingKey = 1;
      if (TILING_KEY_IS(1)) {
          tilingKey = 1;
      } else if (TILING_KEY_IS(2)) {
          tilingKey = 2;
      } else {
          tilingKey = 1;
      }
      op.Init(x, y, z, value, out, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.formerNum,
            tiling_data.formerLength, tiling_data.formertileNum,
            tiling_data.formertileLength, tiling_data.formerlasttileLength,
            tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
            tiling_data.tailtileLength, tiling_data.taillasttileLength,
            tilingKey, tiling_data.datatype);
      op.Process();
    }else{
      ASSERT( "data type is invalid!");
    }

}

#ifndef __CCE_KT_TEST__
// call of kernel function
void addcmul_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* y, uint8_t* z, uint8_t* value, uint8_t* out,
                       uint8_t* workspace, uint8_t* tiling) {
  addcmul<<<blockDim, l2ctrl, stream>>>(x, y, z, value, out, workspace, tiling);
}
#endif


// #include "kernel_operator.h"

// using namespace AscendC;
// constexpr int32_t BUFFER_NUM = 2;
// template<typename T>
// class KernelAddcmul {
//  public:
//   __aicore__ inline KernelAddcmul() {}
//   __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR value, GM_ADDR out,
//                               uint32_t blockLength,
//                               uint32_t tileNum, uint32_t tileLength,
//                               uint32_t lasttileLength, uint32_t formerNum,
//                               uint32_t formerLength, uint32_t formertileNum,
//                               uint32_t formertileLength,
//                               uint32_t formerlasttileLength, uint32_t tailNum,
//                               uint32_t tailLength, uint32_t tailtileNum,
//                               uint32_t tailtileLength,
//                               uint32_t taillasttileLength, uint32_t tilingKey) {
//     ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

//     valueGm.SetGlobalBuffer((__gm__ T*)value , 1);
//     this->value = valueGm.GetValue(0);

//     if (tilingKey == 1) {
//         this->blockLength = blockLength;
//         this->tileNum =
//             tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
//         this->tileLength = tileLength / BUFFER_NUM;
//         this->lasttileLength = lasttileLength;

//         xGm.SetGlobalBuffer((__gm__ T*)x + this->blockLength * GetBlockIdx(),
//                             this->blockLength);
//         yGm.SetGlobalBuffer((__gm__ T*)y + this->blockLength * GetBlockIdx(),
//                             this->blockLength);
//         zGm.SetGlobalBuffer((__gm__ T*)z + this->blockLength * GetBlockIdx(),
//                             this->blockLength);
//         // valueGm.SetGlobalBuffer((__gm__ T*)value + this->blockLength * GetBlockIdx(),
//         //                     this->blockLength);
//         outGm.SetGlobalBuffer(
//             (__gm__ T*)out + this->blockLength * GetBlockIdx(),
//             this->blockLength);
//         }

//     if (tilingKey == 2) {
//       this->formerNum = formerNum;
//       this->formerLength = formerLength;
//       this->formertileNum = formertileNum;
//       this->formertileLength = formertileLength;
//       this->formerlasttileLength = formerlasttileLength;

//       this->tailNum = tailNum;
//       this->tailLength = tailLength;
//       this->tailtileNum = tailtileNum;
//       this->tailtileLength = tailtileLength;
//       this->taillasttileLength = taillasttileLength;

//       if (GetBlockIdx() < this->formerNum) {  //分到大块核的处理
//         this->tileLength = this->formertileLength / BUFFER_NUM;
//         this->lasttileLength = this->formerlasttileLength;
//         this->tileNum = this->formertileNum * BUFFER_NUM;
//         xGm.SetGlobalBuffer(
//             (__gm__ T*)x + this->formerLength * GetBlockIdx(),
//             this->formerLength);
//         yGm.SetGlobalBuffer(
//             (__gm__ T*)y + this->formerLength * GetBlockIdx(),
//             this->formerLength);
//         zGm.SetGlobalBuffer(
//             (__gm__ T*)z + this->formerLength * GetBlockIdx(),
//             this->formerLength);
//         // valueGm.SetGlobalBuffer(
//         //     (__gm__ T*)value + this->formerLength * GetBlockIdx(),
//         //     this->formerLength);
//         outGm.SetGlobalBuffer(
//             (__gm__ T*)out + this->formerLength * GetBlockIdx(),
//             this->formerLength);
//       } else {  //分到小块核的处理，需要处理的数据量比大核少alignNum个
//         this->tileLength = this->tailtileLength / BUFFER_NUM;
//         this->lasttileLength = this->taillasttileLength;
//         this->tileNum = this->tailtileNum * BUFFER_NUM;
//         xGm.SetGlobalBuffer(
//             (__gm__ T*)x + this->formerLength * this->formerNum +
//                 this->tailLength * (GetBlockIdx() - this->formerNum),
//             this->tailLength);
//         yGm.SetGlobalBuffer(
//             (__gm__ T*)y + this->formerLength * this->formerNum +
//                 this->tailLength * (GetBlockIdx() - this->formerNum),
//             this->tailLength);
//         zGm.SetGlobalBuffer(
//             (__gm__ T*)z + this->formerLength * this->formerNum +
//                 this->tailLength * (GetBlockIdx() - this->formerNum),
//             this->tailLength);
//         // valueGm.SetGlobalBuffer(
//         //     (__gm__ T*)value + this->formerLength * this->formerNum +
//         //         this->tailLength * (GetBlockIdx() - this->formerNum),
//         //     this->tailLength);
//         outGm.SetGlobalBuffer(
//             (__gm__ T*)out + this->formerLength * this->formerNum +
//                 this->tailLength * (GetBlockIdx() - this->formerNum),
//             this->tailLength);
//       }
//     }

//     pipe.InitBuffer(inQueueIN, BUFFER_NUM, this->tileLength * 4 * sizeof(T));
//     pipe.InitBuffer(outQueueOUT, BUFFER_NUM, this->tileLength * sizeof(T));
//   }
//   __aicore__ inline void Process() {
//     int32_t loopCount = this->tileNum * BUFFER_NUM;
//     for (int32_t i = 0; i < loopCount; i++) {
//       CopyIn(i);
//       Compute(i);
//       CopyOut(i);
//     }
//   }

//  private:
//   __aicore__ inline void CopyIn(int32_t progress) {
//     LocalTensor<T> inLocal = inQueueIN.AllocTensor<T>();

//     if (BUFFER_NUM == 1) {
//       if (progress == this->tileNum - 1) {
//         if (progress == 0) {
//             //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
//             DataCopy(inLocal[0], xGm[0], this->tileLength);
//             DataCopy(inLocal[this->tileLength], yGm[0], this->tileLength);
//             DataCopy(inLocal[2 * (this->tileLength)], zGm[0], this->tileLength);
//             DataCopy(inLocal[3 * (this->tileLength)], zGm[0], this->tileLength);
//         } else {
//             //将最后一个分块的起始地址向前移动tileLength-lasttileLength
//             DataCopy(
//                 inLocal[0],
//                 xGm[(progress - 1) * this->tileLength + this->lasttileLength],
//                 this->tileLength);
//             DataCopy(
//                 inLocal[this->tileLength],
//                 yGm[(progress - 1) * this->tileLength + this->lasttileLength],
//                 this->tileLength);
//             DataCopy(
//                 inLocal[2 * (this->tileLength)],
//                 zGm[(progress - 1) * this->tileLength + this->lasttileLength],
//                 this->tileLength);
//             // DataCopy(
//             //     inLocal[3 * (this->tileLength)],
//             //     valueGm[(progress - 1) * this->tileLength + this->lasttileLength],
//             //     this->tileLength);
//         }
//       } else {
//             DataCopy(inLocal[0], xGm[progress * this->tileLength],
//                     this->tileLength);
//             DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
//                     this->tileLength);
//             DataCopy(inLocal[2 * (this->tileLength)],
//                     zGm[progress * this->tileLength], this->tileLength);
//             // DataCopy(inLocal[3 * (this->tileLength)],
//             //         valueGm[progress * this->tileLength], this->tileLength);
//       }
//     }
//     if (BUFFER_NUM == 2) {
//       //开启double
//       //buffer时，由于将输入数据分成了相等的2部分，分块大小为不开启double
//       //buffer的一半， 所以需要对最后两个分块数据的起始地址做处理
//       if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
//           (progress == (this->tileNum * BUFFER_NUM - 1))) {
//         //分块大小变为tileLength的一半
//         //倒数第2个分块数据的起始地址向前移动（tileLength-lasttileLength)，最后一个分块的起始地址以此为基础进行移动
//         DataCopy(
//             inLocal[0],
//             xGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
//             (this->tileLength));
//         DataCopy(
//             inLocal[this->tileLength],
//             yGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
//             (this->tileLength));
//         DataCopy(
//             inLocal[2 * (this->tileLength)],
//             zGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
//             (this->tileLength));
//         // DataCopy(
//         //     inLocal[3 * (this->tileLength)],
//         //     valueGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
//         //     (this->tileLength));

//       }

//       else {
//         DataCopy(inLocal[0], xGm[progress * (this->tileLength)],
//                  (this->tileLength));
//         DataCopy(inLocal[this->tileLength], yGm[progress * this->tileLength],
//                  this->tileLength);
//         DataCopy(inLocal[2 * (this->tileLength)],
//                  zGm[progress * this->tileLength], this->tileLength);
//         // DataCopy(inLocal[3 * (this->tileLength)],
//         //          valueGm[progress * this->tileLength], this->tileLength);
//       }
//     }

//     inQueueIN.EnQue(inLocal);
//   }
//   __aicore__ inline void Compute(int32_t progress) {
//     LocalTensor<T> inLocal = inQueueIN.DeQue<T>();
//     LocalTensor<T> xLocal = inLocal;
//     LocalTensor<T> yLocal = inLocal[this->tileLength];
//     LocalTensor<T> zLocal = inLocal[2 * (this->tileLength)];
//     //LocalTensor<T> valueLocal = inLocal[3 * (this->tileLength)];

//     LocalTensor<T> outLocal = outQueueOUT.AllocTensor<T>();

//     Mul(outLocal, yLocal, zLocal, this->tileLength);
//     //Mul(outLocal, outLocal, valueLocal, this->tileLength);
//     Muls(outLocal, outLocal, this->value, this->tileLength);
//     Add(outLocal, xLocal, outLocal, this->tileLength);
//     outQueueOUT.EnQue<T>(outLocal);

//     inQueueIN.FreeTensor(inLocal);
//   }
//   __aicore__ inline void CopyOut(int32_t progress) {
//     LocalTensor<T> outLocal = outQueueOUT.DeQue<T>();

//     if (BUFFER_NUM == 1) {
//       if (progress == this->tileNum - 1) {
//         if (progress == 0) {
//           //如果只有一包，则搬运的起始地址为0，tileLength为实际分块的数据量
//           DataCopy(outGm[0], outLocal, this->tileLength);
//         } else {
//           //将最后一个分块的起始地址向前移动tileLength-lasttileLength
//           DataCopy(
//               outGm[(progress - 1) * this->tileLength + this->lasttileLength],
//               outLocal, this->tileLength);
//         }
//       } else {
//         DataCopy(outGm[progress * this->tileLength], outLocal,
//                  this->tileLength);
//       }
//     }
//     if (BUFFER_NUM == 2) {
//       //开启double
//       //buffer时，由于将输入数据分成了相等的2部分，分块大小为不开启double
//       //buffer的一半， 所以需要对最后两个分块数据的起始地址做处理
//       if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
//           (progress == (this->tileNum * BUFFER_NUM - 1))) {
//         //分块大小变为tileLength的一半
//         //倒数第2个分块数据的起始地址向前移动（tileLength-lasttileLength)，最后一个分块的起始地址以此为基础进行移动
//         DataCopy(
//             outGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
//             outLocal, (this->tileLength));
//       }

//       else {
//         DataCopy(outGm[progress * (this->tileLength)], outLocal,
//                  (this->tileLength));
//       }
//     }

//     outQueueOUT.FreeTensor(outLocal);
//   }

//  private:
//   TPipe pipe;
//   // TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueZ;
//   TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
//   TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
//   GlobalTensor<T> xGm;
//   GlobalTensor<T> yGm;
//   GlobalTensor<T> zGm;
//   GlobalTensor<T> valueGm;
//   GlobalTensor<T> outGm;
//   T value;
//   uint32_t blockLength;
//   uint32_t tileNum;
//   uint32_t tileLength;
//   uint32_t lasttileLength;
//   uint32_t formerNum;
//   uint32_t formerLength;
//   uint32_t formertileNum;
//   uint32_t formertileLength;
//   uint32_t formerlasttileLength;
//   uint32_t tailNum;
//   uint32_t tailLength;
//   uint32_t tailtileNum;
//   uint32_t tailtileLength;
//   uint32_t taillasttileLength;
// };

// extern "C" __global__ __aicore__ void addcmul(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR value, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling) {
//     GET_TILING_DATA(tiling_data, tiling);
//     // TODO: user kernel impl
//     KernelAddcmul<half> op;

//     uint32_t tilingKey = 1;
//     if (TILING_KEY_IS(1)) {
//         tilingKey = 1;
//     } else if (TILING_KEY_IS(2)) {
//         tilingKey = 2;
//     } else {
//         tilingKey = 1;
//     }
//     op.Init(x, y, z, value, out, tiling_data.blockLength,
//           tiling_data.tileNum, tiling_data.tileLength,
//           tiling_data.lasttileLength, tiling_data.formerNum,
//           tiling_data.formerLength, tiling_data.formertileNum,
//           tiling_data.formertileLength, tiling_data.formerlasttileLength,
//           tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
//           tiling_data.tailtileLength, tiling_data.taillasttileLength,
//           tilingKey);
//     op.Process();
// }

// #ifndef __CCE_KT_TEST__
// // call of kernel function
// void addcmul_do(uint32_t blockDim, void* l2ctrl, void* stream,
//                        uint8_t* x, uint8_t* y, uint8_t* z, uint8_t* value, uint8_t* out,
//                        uint8_t* workspace, uint8_t* tiling) {
//   addcmul<<<blockDim, l2ctrl, stream>>>(x, y, z, value, out, workspace, tiling);
// }
// #endif
